#include "graphColoring.h"

//-----------------------used for sort---------------------
struct Color_with_vertexCount {
	int c;
	int vertexCount;
};

struct Sort_by_vertexCount {		// ascending
	bool operator()(Color_with_vertexCount const &a, Color_with_vertexCount const &b) {
		return a.vertexCount < b.vertexCount;
	}
};

struct Sort_by_vertexCountDesc {	// descending
	bool operator()(Color_with_vertexCount const &a, Color_with_vertexCount const &b) {
		return a.vertexCount > b.vertexCount;
	}
};

// obtain the degree of a node
int __device__ degree(int vertex, int *degreeList){
	return degreeList[vertex];
}

// adapted from Grossnet 2011
// saturation of a vertex called from GPU
int __device__ saturation(int vertex, int *adjacencyList, int *graphColors, int maxDegree, int start, int end){
	int saturation = 0;	
	int colors[TEMP_COLOR_LENGTH];			
	for (int j=0; j<TEMP_COLOR_LENGTH; j++)	{
		colors[j] = 0;
	}
	
	for (int i=0; i<maxDegree; i++){
		if (adjacencyList[vertex*maxDegree + i] < start)
			continue;
		
		if (adjacencyList[vertex*maxDegree + i] > end)
			break;
		
		if (adjacencyList[vertex*maxDegree + i] != -1)
			colors[ graphColors[adjacencyList[vertex*maxDegree + i]] ] = 1;			// at each colored set the array to 1
		else
			break;
	}
	
	// count the number of 1's but skip uncolored
	for (int i=1; i<TEMP_COLOR_LENGTH; i++)	{
		if (colors[i] == 1) {
			saturation++;
		}
	}
	
	return saturation;
}


// adapted from Grossnet 2011
// colors the vertex with the min possible color
int __device__ color(int vertex, int *adjacencyList, int *graphColors, int maxDegree, int numColored, int start, int end, int disp){
	int colors[TEMP_COLOR_LENGTH];			
	for (int j=0; j<TEMP_COLOR_LENGTH; j++)	{
		colors[j] = 0;
	}
	
	if (graphColors[vertex] == 0) {
		numColored++;
	}
	
	for (int i=0; i<maxDegree; i++) {						// set the index of the color to 1	
		// Limits color checking to subgraph
		
		if (adjacencyList[vertex*maxDegree + i] != -1)
			colors[  graphColors[  adjacencyList[vertex*maxDegree + i]  ]  ] = 1;
		else 
			break;
	}
	
	
	// nodes still equal to 0 are unassigned
	for (int i=1; i<TEMP_COLOR_LENGTH; i++)		// OPTION2	
	//for (int i=1; i<maxDegree+1; i++)				// OPTION1				
		if (colors[i] != 1){
			if (disp == 0){
				graphColors[vertex] = i;
				break;
			}
			else
				disp--;
		}
	
	return numColored;
}

// asynchronized recoloring 
__global__ void reColor_async(int totalColorsInPreviousIteration, int *orderOfColorsD, int *adjacencyList, int *graphColors, int maxDegree, int numColored, int *startPartitionList, int *endPartitionList, int *orderList, int *previousColors, int *randomListD){
	int randomCount = 0;
	int partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	int start = startPartitionList[partitionIndex];
	int end = endPartitionList[partitionIndex];
	for (int i = 0; i<      totalColorsInPreviousIteration+1; ++i) {        // in case there is color 0: not colored node
		int currColor = orderOfColorsD[i];
		if (currColor == -1) continue;  // error, should not be -1
		for (int j = start; j< end; ++j) {
			randomCount++;
			randomCount = randomCount%10;
			int index = orderList[j];
			if (previousColors[j] == currColor && index != -1)      // match! go color it.
				numColored = color(index,adjacencyList, graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
		}
	}
}

// synchronized recoloring
__global__ void reColor_sync(int totalColorsInPreviousIteration, int *orderOfColorsD, int *adjacencyList, int *graphColors, int maxDegree, int numColored, int *startPartitionList, int *endPartitionList, int *orderList, int *previousColors, int *randomListD, int currColor){
	int randomCount = 0;
	int partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	int start = startPartitionList[partitionIndex];
	int end = endPartitionList[partitionIndex];
		for (int j = start; j< end; ++j) {
			randomCount++;
			randomCount = randomCount%10;
			int index = orderList[j];
			if (previousColors[j] == currColor && index != -1)      // match! go color it.
				numColored = color(index,adjacencyList, graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
		}
}


// coloring using SL vertex ordering
__global__ void colorGraph_SL(int *adjacencyList, int *graphColors, int *degreeList, int sizeGraph, int maxDegree,
					int *startPartitionListD, int *endPartitionListD, int *randomListD, int *inTheGraphList, int *orderList) {
	int start, end, partitionIndex;
	int numColored = 0;
	int randomCount = 0;

	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];

	// get the least degree
	for (int i = end-1; i >= start; --i) { 	//go from the end to the front. 
		int minDegree = maxDegree+1;
		int minVertex = -1;
		for (int j = start; j< end; ++j) {
			if (inTheGraphList[j] == 0) 	//already removed from the graph
				continue;
			int degree = 0;		//degree in subgraph
			for (int k = 0; k < maxDegree; ++k) {
				if ((adjacencyList[j*maxDegree+k] >= start && adjacencyList[j*maxDegree+k] < end) && inTheGraphList[adjacencyList[j*maxDegree+k]] == 1){
				//if (adjacencyList[j*maxDegree+k] > -1 && inTheGraphList[adjacencyList[j*maxDegree+k]] == 1){
					degree++;
				}
			}
			if (degree < minDegree) {
				minDegree = degree;
				minVertex = j;
			}
		}

		//assign the vertex with the least degree to the last slot.
		orderList[i] = minVertex;	// back to the original index.
		inTheGraphList[minVertex] = 0;	//remove from the subgraph.
	}
	
	// color the subgraph following the order
	for (int i=start; i<end; i++) {
		randomCount++;
		randomCount = randomCount%10;
		int index = orderList[i];
		if (index != -1)	// if index = -1, ordering is messed up. 
			numColored = color(index,adjacencyList, graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
	}
}

//coloring using LF ordering
__global__ void colorGraph_LF(int *adjacencyList, int *graphColors, int *degreeList, int sizeGraph, int maxDegree, 
								int *startPartitionListD, int *endPartitionListD, int *randomListD, int *orderedDegreeListD) {
	int start, end, partitionIndex;
	int subGraphSize, numColored = 0;
	int randomCount = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];
        subGraphSize = end - start;
    
        //We use a list which orderes the vertices in the sub-graph in non-increasing degree order
        
	// if first iteration, get the order
	for (int k=start; k<end; k++) {
		orderedDegreeListD[k] = k;    		        
	}
	
	//Using a simple insertion sort to order the list
	int x, y, key;
	for(y = start + 1; y < end; y++) {
	   key = orderedDegreeListD[y];
	   for(x = y - 1; (x >= start) && (degree(orderedDegreeListD[x], degreeList) < degree(key, degreeList)); x--) {
			 orderedDegreeListD[x+1] = orderedDegreeListD[x];
	   }
	   orderedDegreeListD[x+1] = key;    //Put key into its proper location
	}
		
	while (numColored < subGraphSize) {
		randomCount++;
		randomCount = randomCount%10;
		
		for (int i=start; i<end; i++) {
			numColored = color(orderedDegreeListD[i],adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
		}
	}
}


// coloring using SDO vertex ordering
// adapted from Grossnet 2011
__global__ void colorGraph_SDO(int *adjacencyList, int *graphColors, int *degreeList, int sizeGraph, int maxDegree, 
								int *startPartitionListD, int *endPartitionListD, int *randomListD) {
	int start, end, partitionIndex;
	int subGraphSize, numColored = 0;
	int satDegree, max, index;
	int randomCount = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];
	subGraphSize = end - start;
	
	while (numColored < subGraphSize){
		randomCount++;
		randomCount = randomCount%10;
		
		max = -1;
		
		for (int i=start; i<end; i++){
			if (graphColors[i] == 0) {		// not colored
				satDegree = saturation(i,adjacencyList,graphColors, maxDegree, start, end);
				
				if (satDegree > max){
					max = satDegree;
					index = i;				
				}
				
				if (satDegree == max){
					if (degree(i,degreeList) > degree(index,degreeList))
						index = i;
				}			
				
			}
			numColored = color(index,adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
		}
	}
}

//LF followed by SDO
__global__ void colorGraph_LF_SDO(int *adjacencyList, int *graphColors, int *degreeList, int sizeGraph, int maxDegree, 
								int *startPartitionListD, int *endPartitionListD, int *randomListD, int *LFSDODegreeListD) {
	int start, end, partitionIndex;
	int subGraphSize, numColored = 0;
	int randomCount = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];
	subGraphSize = end - start;

        //Just as in LF, We use a list which orderes the vertices in the sub-graph in non-increasing degree order
        
	// if first iteration, get the order
	for (int k=start; k<end; k++){
		LFSDODegreeListD[k] = k;    		        
	}
	
	//Using a simple insertion sort to order the list
	int x, y, key;
	for(y = start + 1; y < end; y++) {
		key = LFSDODegreeListD[y];
		for(x = y - 1; (x >= start) && (degree(LFSDODegreeListD[x], degreeList) <= degree(key, degreeList)); x--){
			if(degree(LFSDODegreeListD[x], degreeList) == degree(key, degreeList)){
			//Thus there is tie when only degree is cosidered. To break the tie we need to consider saturation
				if(saturation(LFSDODegreeListD[x],adjacencyList,graphColors, maxDegree, start, end) >= saturation(key,adjacencyList,graphColors, maxDegree, start, end)){
					continue;
				}
			}
			LFSDODegreeListD[x+1] = LFSDODegreeListD[x];
	   	}
		LFSDODegreeListD[x+1] = key;    //Put key into its proper location
	}
			
	while (numColored < subGraphSize){
		randomCount++;
		randomCount = randomCount%10;
		
		for (int i=start; i<end; i++){
			numColored = color(LFSDODegreeListD[i],adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
		}
	}
	
}



//solve conflicts using SL ordering
__global__ void conflictSolveSL(int *adjacencyList, int *conflict, int *graphColors, int *degreeList,
		int sizeGraph, int maxDegree, int *startPartitionListD, int *endPartitionListD, int *randomListD, int *inTheGraphList, int *orderList){
	
	int start, end, partitionIndex;
	int numColored = 0;
	int randomCount = 0;

	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];

	for (int i=start; i<end; i++){
		randomCount++;
		randomCount = randomCount%10;
		int index = orderList[i];
		if (index != -1 && graphColors[index] == 0)	// if index = -1, ordering is messed up. 
			numColored = color(index,adjacencyList, graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
	}
}

//solve conflicts using LF ordering
__global__ void conflictSolveLF(int *adjacencyList, int *conflict, int *graphColors, int *degreeList, 
								int sizeGraph, int maxDegree, int *startPartitionListD, int *endPartitionListD, int *randomListD, int *orderedDegreeListD){
	int start, end, partitionIndex;
	int numColored = 0;
	int randomCount = 0;
	int numOfInitialConflicts = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];

	// Count the number of conflicts
	for (int i=start; i<end; i++) {
		if (graphColors[i] == 0)
			numOfInitialConflicts++;
	}
	
	while (numOfInitialConflicts > 0){
		randomCount++;
		randomCount = randomCount%10;
	
		for (int i=start; i<end; i++){
			if (graphColors[orderedDegreeListD[i]] == 0){   //Not colored
	  			numColored = color(orderedDegreeListD[i],adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
				numOfInitialConflicts--;
			}
		}
	}
}

//solve conflicts using LF ordering then SDO 
__global__ void conflictSolveLFSDO(int *adjacencyList, int *conflict, int *graphColors, int *degreeList, 
								int sizeGraph, int maxDegree, int *startPartitionListD, int *endPartitionListD, int *randomListD, int *LFSDODegreeListD){
	int start, end, partitionIndex;
	int numColored = 0;
	int randomCount = 0;
	int numOfInitialConflicts = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];

	// Count the number of conflicts
	for (int i=start; i<end; i++) {
		if (graphColors[i] == 0)
			numOfInitialConflicts++;
	}
	
	while (numOfInitialConflicts > 0){
		randomCount++;
		randomCount = randomCount%10;
	
		for (int i=start; i<end; i++){
			if (graphColors[LFSDODegreeListD[i]] == 0){   //Not colored
				numColored = color(LFSDODegreeListD[i],adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
				numOfInitialConflicts--;
			}
		}
	}
}


// solve the conflicts using SDO 
// adapted from Grossnet 2011
//recolors nodes where we have conflicts (conflicts in a subgraph, no communication between subgraphs)
__global__ void conflictSolveSDO(int *adjacencyList, int *conflict, int *graphColors, int *degreeList, 
								int sizeGraph, int maxDegree, int *startPartitionListD, int *endPartitionListD, int *randomListD){
	int start, end, index, partitionIndex;
	int numColored = 0;
	int satDegree, max;
	int randomCount = 0;
	int numOfInitialConflicts = 0;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];
	
	// Count the number of conflicts
	for (int i=start; i<end; i++) {
		if (graphColors[i] == 0)
			numOfInitialConflicts++;
    }
	
    while (numOfInitialConflicts > 0){
        max = -1;
        randomCount++;
		randomCount = randomCount%10;
        
        for (int i=start; i<end; i++){
            if (graphColors[i] == 0)                        // not colored
            {
				satDegree = saturation(i,adjacencyList,graphColors, maxDegree, start, end);
				
                if (satDegree > max){
                    max = satDegree;
                    index = i;                              
                }
				
                if (satDegree == max){
                    if (degree(i,degreeList) > degree(index,degreeList))
                        index = i;
                }
				
				numColored = color(index,adjacencyList,graphColors, maxDegree, numColored, start, end, randomListD[partitionIndex*10 + randomCount]);
				
				if (i == index)
					numOfInitialConflicts--;
            }
        }
    }
}

//adapted from Grossnet 2011
//color the graph using natural ordering. 
__global__ void colorGraph_FF(int *adjacencyListD, int *colors, int size, int maxDegree, int *startPartitionListD, int *endPartitionListD){
	int i, j, start, end, partitionIndex;
	int numColors = 0;
	
	//int subGraphSize;
	//subGraphSize = size/(gridDim.x * blockDim.x);
	//start = (size/gridDim.x * blockIdx.x) + (subGraphSize * threadIdx.x);
	//end = start + subGraphSize;
	
	partitionIndex = (blockIdx.x * blockDim.x) + threadIdx.x;
	start = startPartitionListD[partitionIndex];
	end = endPartitionListD[partitionIndex];
	
	
	int degreeArray[TEMP_COLOR_LENGTH];
	for (i=start; i<end; i++)
	{
		for(j=0; j<maxDegree; j++)
			degreeArray[j] = j+1;
		
		
		for (j=0; j<maxDegree; j++){
			int vertexNeigh = i*maxDegree + j;
			
			if (adjacencyListD[vertexNeigh] != -1){
				if (colors[ adjacencyListD[vertexNeigh] ] != 0)
					degreeArray[ colors[adjacencyListD[vertexNeigh]] -1 ] = 0;
			}
			else
				break;
		}
		
		
		for (j=0; j<maxDegree; j++)
			if(degreeArray[j] != 0){
				colors[i] = degreeArray[j];
				break;
			}
		
		if(colors[i] > numColors)
			numColors = colors[i];		
	}
}



//----------------------- Detects conflicts -----------------------//
// adapted from adapted from Grossnet 2011
// each thread deals with 1 vertex from boundary list
// 		set the conflicted color to 0
// 		set its value in the conflict list to point to the node
__global__ void conflictsDetection(int *adjacentListD, int *boundaryListD, int *colors, int *conflictD, long size, int boundarySize, int maxDegree){
	
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	int nodeFrom, nodeTo;
	
	if (idx < boundarySize){
		nodeFrom = boundaryListD[idx];		

		for (int i=0; i<maxDegree; i++)
		{
			nodeTo = adjacentListD[nodeFrom*maxDegree + i];
			
			if (nodeTo == -1)
				break;
			
			if (nodeFrom>=nodeTo && (colors[nodeFrom] == colors[nodeTo]))
			{
				conflictD[idx] = nodeFrom;	
				colors[nodeFrom] = 0;				
			}		
		}
	}
}


//----------------------- Main -----------------------//
// ordering: 0=FF, 1=SDO, 2=SL, 3=LF. 4=LFSDO
extern "C"
void cudaGraphColoring(int *adjacentList, int *boundaryList, int *graphColors, int *degreeList, int *conflict, int boundarySize, 
						int maxDegree, int graphSize, int passes, int subsizeBoundary, int _gridSize, int _blockSize, 
						int *startPartitionList, int *endPartitionList, int *randomList, int numRand, int ordering, 
						int recoloring, int sync, int recolor_order) {
	
	int *adjacentListD, *colorsD, *boundaryListD, *degreeListD, *conflictListD, *startPartitionListD, *endPartitionListD, *randomListD;     
	int *inTheGraphListD, *orderListD, *previousColorsD, *orderOfColorsD;	// for SL
	int *orderedDegreeListD; // for LF
	int *LFSDODegreeListD; // for LFSDO
	
	int *orderOfColors = new int[TEMP_COLOR_LENGTH];
	
	//store best color
	int bestColor = TEMP_COLOR_LENGTH;
	int *bestColors = new int[graphSize];
	
	int gridsize = ceil((float)boundarySize/(float)(CONFLICT_BLOCK_SIZE));
	int blocksize = CONFLICT_BLOCK_SIZE;
	int *numConflicts;
	
	
	cudaEvent_t start_col, start_confl, stop_col, stop_confl, start_mem, stop_mem;         
        float elapsedTime_memory, elapsedTime_col, elapsedTime_confl; 

	
	
	//-------------- memory transfer -----------------!
	cudaEventCreate(&start_mem); 
        cudaEventCreate(&stop_mem); 
        cudaEventRecord(start_mem, 0); 
	
	
	cudaMalloc((void**)&adjacentListD, graphSize*maxDegree*sizeof(int));
	cudaMalloc((void**)&colorsD, graphSize*sizeof(int));
	cudaMalloc((void**)&boundaryListD, boundarySize*sizeof(int));
	cudaMalloc((void**)&degreeListD, graphSize*sizeof(int));
	cudaMalloc((void**)&numConflicts, 1*sizeof(int));
	cudaMalloc((void**)&conflictListD, boundarySize*sizeof(int));
	cudaMalloc((void**)&startPartitionListD, _gridSize*_blockSize*sizeof(int));
	cudaMalloc((void**)&endPartitionListD, _gridSize*_blockSize*sizeof(int));
	cudaMalloc((void**)&randomListD, numRand*sizeof(int));
	
	//for SL only
	if (ordering == 2) {
		cudaMalloc((void**)&inTheGraphListD, graphSize*sizeof(int));
		cudaMalloc((void**)&orderListD, graphSize*sizeof(int));
	}
	
	// for LF only
	if (ordering == 3) {
	    cudaMalloc((void**)&orderedDegreeListD, graphSize*sizeof(int));  //A list of vertices in non-decreasing order of degree
	}

         // for LFSDO only
	if (ordering == 4) {
        cout<<"creating LFSDODegreeListD\n";
	    cudaMalloc((void**)&LFSDODegreeListD, graphSize*sizeof(int));  //A list of vertices in non-decreasing order of degree with ties based on saturation
	}	

	// for recoloring
	cudaMalloc((void**)&previousColorsD, graphSize*sizeof(int));		// store previous colors
	cudaMalloc((void**)&orderOfColorsD, TEMP_COLOR_LENGTH*sizeof(int));	//store the recolor order
	
	
	cudaMemcpy(adjacentListD, adjacentList, graphSize*maxDegree*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(colorsD, graphColors, graphSize*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(boundaryListD, boundaryList, boundarySize*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(degreeListD, degreeList, graphSize*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(startPartitionListD, startPartitionList, _gridSize*_blockSize*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(endPartitionListD, endPartitionList, _gridSize*_blockSize*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(randomListD, randomList, numRand*sizeof(int), cudaMemcpyHostToDevice);
	
	//for SL only
	if (ordering == 2) {
		cudaMemset(inTheGraphListD, 1, graphSize*sizeof(int));	//1 for in the graph, 0 for not in the graph
		cudaMemset(orderListD, -1, graphSize*sizeof(int));
	}
	
	// for LF only
	if (ordering == 3) {
		cudaMemcpy(orderedDegreeListD, degreeList, graphSize*sizeof(int), cudaMemcpyHostToDevice);
	}

      	// for LFSDO only
	if (ordering == 4) {
		cudaMemcpy(LFSDODegreeListD, degreeList, graphSize*sizeof(int), cudaMemcpyHostToDevice);
	}
        

	// for recoloring
	cudaMemcpy(previousColorsD, graphColors, graphSize*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemset(orderOfColorsD, -1, TEMP_COLOR_LENGTH*sizeof(int));

	cudaEventRecord(stop_mem, 0); 
        cudaEventSynchronize(stop_mem); 
	
	
	dim3 dimGrid_col(_gridSize);
	dim3 dimBlock_col(_blockSize);
	
	dim3 dimGrid_confl(gridsize);
	dim3 dimBlock_confl(blocksize);
	
//!-----------------recoloring iteration------------------------!
//Note: The following pieace of code was written to automate all our experiments (in order to save time/effort as well as
//in order to deal with huge traffic on the server in days running to final presentation
//As a result, it has become a bit messy despite our best efforts. However, it does not contain any new/intresting idea/implementation
//so the reader can safely skip this portion

	int totalColorsInPreviousIteration = 0;
	
	for (int recolor_it = 0; recolor_it <= recoloring; ++recolor_it) {	// do the whole process multiple times
		cout << endl << "recoloring iteration: " << recolor_it << endl;
		
		//-------------- Sequential Graph coloring -----------------!
		
		cudaEventCreate(&start_col); 
		cudaEventCreate(&stop_col); 
		cudaEventRecord(start_col, 0); 
        
          if(recolor_it == 0){	// the first coloring iteration. 
	
			switch (ordering) {
				case 0:
					//colorGraph_FF<<<dimGrid_col, dimBlock_col>>>(adjacentListD, colorsD, graphSize, maxDegree);				// First Fit
					break;
				case 1:
					colorGraph_SDO<<<dimGrid_col, dimBlock_col>>>(adjacentListD, colorsD, degreeListD,
													graphSize, maxDegree, startPartitionListD, endPartitionListD, randomListD);		// SDO improved
					break;
				case 2:
					colorGraph_SL<<<dimGrid_col, dimBlock_col>>>(adjacentListD, colorsD, degreeListD,
									graphSize, maxDegree, startPartitionListD, endPartitionListD, randomListD, 
									inTheGraphListD, orderListD);		//SL ordering
					break;
				case 3:
					colorGraph_LF<<<dimGrid_col, dimBlock_col>>>(adjacentListD, colorsD, degreeListD, 
									graphSize, maxDegree, startPartitionListD, endPartitionListD, randomListD, 
									orderedDegreeListD);	//LF ordering
					break;
							case 4:
									
					colorGraph_LF_SDO<<<dimGrid_col, dimBlock_col>>>(adjacentListD, colorsD, degreeListD, 
									graphSize, maxDegree, startPartitionListD, endPartitionListD, randomListD, 
									LFSDODegreeListD);	//LFSDO ordering
					break;
	
				default:
					return;
			}
          }	// end of recolor_it == 0
          else{  //recoloring 
			if(sync == 0){ //recor async
				if(ordering == 2) { //SL
					reColor_async<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, orderListD, previousColorsD, randomListD);
				}
				else if(ordering == 3){ //LF
					reColor_async<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, orderedDegreeListD, previousColorsD, randomListD);
				}
				else if(ordering == 4){ //LFSDO
					reColor_async<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, LFSDODegreeListD, previousColorsD, randomListD);
				}
                        }	// end of recolor async
                        else{ //recolor synch
		   		if(ordering == 2) { //SL
					for(int i = 0; i < totalColorsInPreviousIteration+1; ++i){
						reColor_sync<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, orderListD, previousColorsD, randomListD,orderOfColors[i]);
						cudaMemset(conflictListD, -1, boundarySize*sizeof(int));
						conflictsDetection<<<dimGrid_confl, dimBlock_confl>>>(adjacentListD, boundaryListD, colorsD, conflictListD, graphSize, boundarySize, maxDegree);
						conflictSolveSL<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize,
																maxDegree, startPartitionListD, endPartitionListD, randomListD, inTheGraphListD, orderListD);
					}
                }
                else if(ordering ==3) { //LF
					for(int i = 0; i < totalColorsInPreviousIteration+1; ++i){
						reColor_sync<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, orderedDegreeListD, previousColorsD, randomListD, orderOfColors[i]);
						cudaMemset(conflictListD, -1, boundarySize*sizeof(int));
						conflictsDetection<<<dimGrid_confl, dimBlock_confl>>>(adjacentListD, boundaryListD, colorsD, conflictListD, graphSize, boundarySize, maxDegree);
						conflictSolveLF<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize, 
																maxDegree, startPartitionListD, endPartitionListD, randomListD, orderedDegreeListD);
					}
                }
				else if(ordering ==4) { //LFSDO
					for(int i = 0; i < totalColorsInPreviousIteration+1; ++i){
			   			reColor_sync<<<dimGrid_col, dimBlock_col>>>(totalColorsInPreviousIteration, orderOfColorsD, adjacentListD, colorsD, maxDegree, 0, startPartitionListD, endPartitionListD, LFSDODegreeListD, previousColorsD, randomListD, orderOfColors[i]);
		           		cudaMemset(conflictListD, -1, boundarySize*sizeof(int));
			   			conflictsDetection<<<dimGrid_confl, dimBlock_confl>>>(adjacentListD, boundaryListD, colorsD, conflictListD, graphSize, boundarySize, maxDegree);
	        	   		conflictSolveLFSDO<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize, 
														maxDegree, startPartitionListD, endPartitionListD, randomListD, LFSDODegreeListD);
                 	}
                }
            }	// end of recolor synch
		}	//end of recoloring
		
		cudaEventRecord(stop_col, 0); 
		cudaEventSynchronize(stop_col); 
		
		cudaEventElapsedTime(&elapsedTime_col, start_col, stop_col);
		cout << "initial coloring time: " << elapsedTime_col << " ms"<< endl;	
				
		
		//-------------- Conflict resolution for multiple passes-----------------!	
		for (int times=1; times<passes; times++){
			
			cudaEventCreate(&start_confl); 
			cudaEventCreate(&stop_confl); 
			cudaEventRecord(start_confl, 0); 
			
			cudaMemset(conflictListD, -1, boundarySize*sizeof(int));
			conflictsDetection<<<dimGrid_confl, dimBlock_confl>>>(adjacentListD, boundaryListD, colorsD, conflictListD, graphSize, boundarySize, maxDegree);
			
			cudaEventRecord(stop_confl, 0); 
			cudaEventSynchronize(stop_confl); 
			
			
			cudaEventCreate(&start_col); 
			cudaEventCreate(&stop_col); 
			cudaEventRecord(start_col, 0); 
		
			switch (ordering) {
				case 0: 
					break;
				case 1:
					conflictSolveSDO<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize, 
															maxDegree, startPartitionListD, endPartitionListD, randomListD);
					break;
				case 2:
					conflictSolveSL<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize,
															maxDegree, startPartitionListD, endPartitionListD, randomListD, inTheGraphListD, orderListD);
					break;
				case 3:
					conflictSolveLF<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize, 
														maxDegree, startPartitionListD, endPartitionListD, randomListD, orderedDegreeListD);
					break;
                                case 4:
					conflictSolveLFSDO<<<dimGrid_col, dimBlock_col>>>(adjacentListD, conflictListD, colorsD, degreeListD, graphSize, 
														maxDegree, startPartitionListD, endPartitionListD, randomListD, LFSDODegreeListD);
					break;
				default:
					return;
			}	
	
			cudaEventRecord(stop_col, 0); 
			cudaEventSynchronize(stop_col); 
			
			cudaEventElapsedTime(&elapsedTime_col, start_col, stop_col); 
			cudaEventElapsedTime(&elapsedTime_confl, start_confl, stop_confl); 
			cout << "Coloring: " << elapsedTime_col << " ms    Conflict: " << elapsedTime_confl << " ms" << endl;
		}	// end of for times 
		
		// if it is the first iteration, we always get the best color
		if (recolor_it == 0) {
			cudaMemcpy(bestColors, colorsD, graphSize*sizeof(int), cudaMemcpyDeviceToHost);
			bestColor = 0;
			for (int i = 0; i< graphSize; ++i) {
				if (bestColors[i] > bestColor)
					bestColor++;
			}
		} 
		
		// if we need recoloring
		if (recoloring > 0) {
			// copy the colorsD to graphColors, get the total number of colors in CPU, and set the as previousColor.
			cudaMemcpy(graphColors, colorsD, graphSize*sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpy(previousColorsD, graphColors, graphSize*sizeof(int), cudaMemcpyHostToDevice);
			// reset the colors
			cudaMemset(colorsD, 0, graphSize*sizeof(int));
	
			totalColorsInPreviousIteration = 0;
		
            for (int i = 0; i< graphSize; ++i) {
				if (graphColors[i] > totalColorsInPreviousIteration)
					totalColorsInPreviousIteration++;
			}
	
			// generate order of colors based on Non-decreasing number of vertices. 
			memset(orderOfColors, -1, TEMP_COLOR_LENGTH*sizeof(int));	//colors start from 1. 
			if(recolor_order > 0){
				vector<Color_with_vertexCount> colors_order;
				for (int i = 0; i< TEMP_COLOR_LENGTH; ++i){
					Color_with_vertexCount color;
					color.c = 0; 
					color.vertexCount = 0;
					colors_order.push_back(color);
				}
			
				for (int i= 0; i< graphSize; ++i)
					colors_order[graphColors[i]].vertexCount++;
				
				if (recolor_it > 0 && totalColorsInPreviousIteration < bestColor) {
					cout << "previous best: " << bestColor;
					cout << "new best scheme" << endl;
					bestColor = totalColorsInPreviousIteration;
					memcpy(bestColors, graphColors, graphSize*sizeof(int));
				}	
			
				if (recolor_order ==1) { //increasing
				// sort the orderOfColors by acsending order of vertexCountForColors.
				   sort (colors_order.begin(), colors_order.begin()+totalColorsInPreviousIteration+1, Sort_by_vertexCount());
				}
				
				else if(recolor_order ==2){ //decreasing
				// sort from high to low
				   sort (colors_order.begin(), colors_order.begin()+totalColorsInPreviousIteration+1, Sort_by_vertexCountDesc());
				}
					
				for (int i = 0; i< totalColorsInPreviousIteration+1; ++i) {
					orderOfColors[i] = colors_order[i].c; 	
				}				
			}	// end if recolor_order > 0
			else {		// reverse order
				for (int i = 0; i< totalColorsInPreviousIteration+1; ++i)
					orderOfColors[i] = totalColorsInPreviousIteration - i;
			}
		    
			cudaMemcpy(orderOfColorsD, orderOfColors, TEMP_COLOR_LENGTH*sizeof(int), cudaMemcpyHostToDevice);
			
			cout <<"total colors: " << totalColorsInPreviousIteration << endl;	
		}	// end of if recoloring > 0

	}	// end of for recolor_it
	
	
	cudaEventCreate(&start_confl); 
        cudaEventCreate(&stop_confl); 
        cudaEventRecord(start_confl, 0); 
	
	cudaMemset(conflictListD, -1, boundarySize*sizeof(int));
	
	cudaMemcpy(colorsD, bestColors, graphSize*sizeof(int), cudaMemcpyHostToDevice);
	cout << "detect conflicts in GPU" << endl;
	
	conflictsDetection<<<dimGrid_confl, dimBlock_confl>>>(adjacentListD, boundaryListD, colorsD, conflictListD, graphSize, boundarySize, maxDegree);
	
	cudaEventRecord(stop_confl, 0); 
        cudaEventSynchronize(stop_confl); 
	
	
	
	//-------------- Cleanup -----------------!
	cudaMemcpy(graphColors, colorsD, graphSize*sizeof(int), cudaMemcpyDeviceToHost);
	cudaMemcpy(conflict, conflictListD, boundarySize*sizeof(int), cudaMemcpyDeviceToHost);
	
	cudaEventElapsedTime(&elapsedTime_memory, start_mem, stop_mem); 
	cudaEventElapsedTime(&elapsedTime_col, start_col, stop_col); 
	cudaEventElapsedTime(&elapsedTime_confl, start_confl, stop_confl); 
	
	cout << endl << "GPU timings ~ Memory transfer: " << elapsedTime_memory  << " ms     Coloring: " 
	<< elapsedTime_col << " ms    Conflict: " << elapsedTime_confl << " ms" << endl; 
	
	
	cudaFree(adjacentListD);
	cudaFree(colorsD);
	cudaFree(boundaryListD);
	cudaFree(degreeListD);
	cudaFree(numConflicts);
	cudaFree(conflictListD);
	cudaFree(startPartitionListD);
	cudaFree(endPartitionListD);
	cudaFree(randomListD);

	//for SL only
	if (ordering == 2) {
		cudaFree(inTheGraphListD);
		cudaFree(orderListD);
	}
	
	//for LF only
	if (ordering == 3) {
		cudaFree(orderedDegreeListD);
	}

        //for LFSDO only
	if (ordering == 4) {
		cudaFree(LFSDODegreeListD);
	}
	
	delete []orderOfColors;
	delete []bestColors;
}

//End of file
